//----------------------------------------------------------------------------
// M A R I T I M E  R E S E A R C H  I N S T I T U T E  N E T H E R L A N D S
//----------------------------------------------------------------------------
// Copyright (C) 2012 - MARIN - All rights reserved - http://www.marin.nl
//----------------------------------------------------------------------------
// Program    : mermaid
// Module     : vbm
// File       : solv.h
// Author     : M.A. de Jong
//----------------------------------------------------------------------------

#ifndef SOLV_H
#define	SOLV_H

#include "kernels/solv_kernels.h"

namespace rrbsolver
{

template <class T>
inline void solv_phase1v1(T* dst2, T* dst1, const T* src, 
    const Grid grid1, const Grid grid0)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, src, &desc, 
      grid0.nx, grid0.ny, grid0.nx * sizeof(T)), 
        "rrbsolver::bind texture @ kernel_solv1v1");
#endif

    dimGrid.x = grid0.cx / NUM_THREADS_X;
    dimGrid.y = grid0.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_solv1v1<T> <<< dimGrid, dimBlock >>> (dst2, dst1, src, 
                                                       grid1, grid0);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_solv1v1");


#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture @ kernel_solv1v1");
#endif
}

template <class T>
inline void solv_phase1v2(T* dst, const T* src, 
    const Grid grid1, const Grid grid0)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, src, &desc, 
      grid0.nx, grid0.ny, grid0.nx * sizeof(T)), 
        "rrbsolver::bind texture @ kernel_solv1v2");
#endif

    dimGrid.x = grid0.cx / NUM_THREADS_X;
    dimGrid.y = grid0.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_solv1v2<T> <<< dimGrid, dimBlock >>> (dst, src, grid1, grid0);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_solv1v2");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture @ kernel_solv1v2");
#endif
}

template <class T>
inline void solv_phase1v3(T* dst, const Grid grid)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, dst, &desc, 
      grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture @ kernel_solv1v3");
#endif

    dimGrid.x = grid.cx / NUM_THREADS_X;
    dimGrid.y = grid.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_solv1v3<T> <<< dimGrid, dimBlock >>> (dst, grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_solv1v3");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture @ kernel_solv1v3");
#endif
}

template <class T>
inline void solv_phase2(T* dst, const Grid grid)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, dst, &desc, 
      grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture @ kernel_solv2");
#endif

    dimGrid.x = grid.cx / NUM_THREADS_X;
    dimGrid.y = grid.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_solv2<T> <<< dimGrid, dimBlock >>> (dst, grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_solv2");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture @ kernel_solv2");
#endif
}

template <class T>
inline void solv_phase3(T* dst, const T* src, 
    const Grid grid1, const Grid grid0)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, src, &desc, 
      grid0.nx, grid0.ny, grid0.nx * sizeof(T)),
        "rrbsolver::bind texture @ kernel_solv3");
#endif

    dimGrid.x = grid0.cx / NUM_THREADS_X;
    dimGrid.y = grid0.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_solv3<T> <<< dimGrid, dimBlock >>> (dst, src, grid1, grid0);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_solv3");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture @ kernel_solv3");
#endif
}

template <class T>
inline void solv_phase4(T* dst, const Grid grid)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, dst, &desc, 
      grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture @ kernel_solv4");
#endif

    dimGrid.x = grid.cx / NUM_THREADS_X;
    dimGrid.y = grid.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_solv4<T> <<< dimGrid, dimBlock >>> (dst, grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_solv4");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture @ kernel_solv4");
#endif
}

template <class T>
inline void solv_phase5(T* dst, const Grid grid)
{
    dimGrid.x = grid.cx / NUM_THREADS_X;
    dimGrid.y = grid.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_solv5<T> <<< dimGrid, dimBlock >>> (dst, grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_solv5");
}

template <class T>
inline void solv_phase6(T* dst, const Grid grid)
{
    dimGrid.x = grid.cx / NUM_THREADS_X;
    dimGrid.y = grid.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_solv6<T> <<< dimGrid, dimBlock >>> (dst, grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_solv6");
}

template <class T>
inline void solv_final(T* dst, const Grid grid)
{
    dimGrid.x = 1;
    dimGrid.y = 1;
    dimBlock.x = DIM_COMPUTE_BLOCK;
    dimBlock.y = DIM_COMPUTE_BLOCK / 2;

    kernel_solvfinal<T> <<< dimGrid, dimBlock >>> (dst, grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_solvfinal");
}

} // namespace rrbsolver

#endif // SOLV_H

